#include <stdio.h>
#include <cuda_runtime.h>

#define N 1024

__global__ void kernel(int *a) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        a[idx] *= 2;  // GPU直接操作主机内存
    }
}

int main() {
    int *a;
    int size = N * sizeof(int);
    
    // 1. 分配零拷贝内存（主机端pinned且映射到设备端）
    cudaHostAlloc(&a, size, cudaHostAllocMapped);
    
    // 初始化数据
    for (int i = 0; i < N; i++) {
        a[i] = i;
    }
    
    // 2. 获取设备端指针
    int *d_a;
    cudaHostGetDevicePointer(&d_a, a, 0);
    
    // 3. 启动内核
    kernel<<<(N + 255)/256, 256>>>(d_a);
    cudaDeviceSynchronize();
    
    // 验证结果
    for (int i = 0; i < 10; i++) {
        printf("%d ", a[i]);
    }
    printf("\n");
    
    // 4. 释放内存
    cudaFreeHost(a);
    
    return 0;
}